CS►parallel computing-高性能计算►-cuda

cuda编程模型

基础

unified memory architecture

accessible from cpu and GPU

编程模型

Kernels

CUDA C通过允许程序员定义C函数(称为内核)来扩展C，当调用时，由N个不同的CUDA线程并行执行N次，而不是像常规的C函数那样只执行一次。

kernel function用global关键字标记，这样的函数能够在GPU上并行计算。

使用global声明说明符定义内核，并使用新的<<<…执行配置语法(参见C语言扩展)。执行内核的每个线程都有一个惟一的线程ID，该ID可以通过内置的threadIdx变量在内核中访问。

以下是两个N维向量的加法，$ C = A + B $:

传统写法：

1
// todo

cuda编程

1
2
3
4
5
6
7
8
9
10
11
12
13
14
// 首先定义内核函数 (采用__global__关键字)
__global__ void VecAdd(float* A, float* B, float* C)
{
    int i = threadIdx.x;
    C[i] = A[i] + B[i];
}

int main()
{
    ...
    // Kernel invocation with N threads
    VecAdd<<<1, N>>>(A, B, C);
    ...
}

这里，N个线程同时计算VecAdd，进行元素加和。

疑问：
并行度设置为1呢，设置为其他数值呢？

效果对比：

源码下载：

线程层次结构

在cuda中，线程ID（threadIdx)是一个三维向量，因此可用一维、二维或三维线程索引来标识线程，从而形成一维、二维或三维线程块(thread block)。这方便我们进行向量、矩阵或卷的计算。

线程的索引及其线程ID以一种简单的方式相互关联:

  • 一维块，线程的索引 = 线程ID
  • 大小是$(D_x, D_y)$的二维块，索引为$(x, y)$的线程对应的线程ID是$(x + y Dx)$
  • 大小是$(Dx, Dy, Dz)$的三维块，索引为$(x, y, z)$的线程对应的线程ID是$(x + y Dx + z Dx Dy)$

例如，下面的代码将大小为$N \times N$的两个矩阵相加，$C = A + B$ :

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
// Kernel definition
__global__ void MatAdd(float A[N][N], float B[N][N],
                       float C[N][N])
{
    int i = threadIdx.x;  // i,j 对应矩阵的两个维度
    int j = threadIdx.y;
    C[i][j] = A[i][j] + B[i][j];
}

int main()
{
    ...
    // Kernel invocation with one block of N * N * 1 threads
    int numBlocks = 1;   // 为什么numBlocks=1？
    dim3 threadsPerBlock(N, N);  
    MatAdd<<<numBlocks, threadsPerBlock>>>(A, B, C);  // 这里采用
    ...
}

FAQ

内核是函数？

是的，kernel is function，能够并行的function。

##

扩展阅读